
#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable
#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable

// 9. source vector is accesed as uint4 vector

__kernel void minMaxP( __global const uint4 *src,
         __global uint *gresult,
         __local uint *lmin,
         __local uint *lmax,
         uint nitems,
	__global uint * gdbg )
{
 // 10. Set up __global memory access pattern.

 uint global_size = get_global_size(0);
 uint count = ( nitems / 4 ) / global_size;
 uint idx =  get_global_id(0);
 uint stride =  global_size;
 uint pmin = (uint) -1;
 uint pmax = 0;

 // 11. First, compute private min, for this work-item.

 for( int n=0; n < count; n++, idx += stride )
 {
  pmin = min( pmin, src[idx].x );
  pmin = min( pmin, src[idx].y );
  pmin = min( pmin, src[idx].z );
  pmin = min( pmin, src[idx].w );

  pmax = max( pmax, src[idx].x );
  pmax = max( pmax, src[idx].y );
  pmax = max( pmax, src[idx].z );
  pmax = max( pmax, src[idx].w );
 }

 // 12. Reduce min values inside work-group.

 if( get_local_id(0) == 0 ){
  lmin[0] = (uint) -1;
  lmax[0] = 0;
 }

 barrier( CLK_LOCAL_MEM_FENCE );

 (void) atom_min( lmin, pmin );
 (void) atom_max( lmax, pmax );

 barrier( CLK_LOCAL_MEM_FENCE );

 // Write out to __global.

 if( get_local_id(0) == 0 ){
  uint ind = 2 * get_group_id(0);
  gresult[ ind ] = lmin[0];
  gresult[ ind + 1 ] = lmax[0];
 }

if( get_global_id(0) == 0 )
{ 
 gdbg[0] = get_num_groups(0);
 gdbg[1] = get_global_size(0);
 gdbg[2] = count;
 gdbg[3] = stride;
} 

}



